Repository navigation
Conversation
Split each pool and sequence row into bounded chunks so multiple CTAs can copy long block-offset tables. Preserve row strides and source layout, and cover the native launch geometry with a regression test. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
Leave 1 GiB outside the caching allocator cap for driver and KV transfer allocations created after executor setup. Cover the resulting allocator fraction with a CPU regression test. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
…ention With an FP8 KV cache, the FP8 context MLA FMHA reads FP8 Q, K and V, and the attention op copies BF16 inputs to FP8 on every call. Let the caller hand it E4M3 K / V directly, and optionally an FP8 Q through quant_q_buffer with quant_scale_qkv, so that a caller that already produces FP8 tensors skips that copy. BF16 inputs keep the existing path. Expose has_fp8_context_mla on TrtllmAttention so that callers can tell when the op runs this FMHA. Cover bit-identical outputs for FP8 and BF16 inputs on the cached-KV context path, and the dtype check. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
Append drafter layers to the target cache manager and propagate their cache layout through disaggregated transfer. Preserve token-parity lifecycle windows and per-rank recurrent state, and validate supported page sizes. Drafters declare what they read through DFlashForCausalLM hooks (reads_managed_ctx_pages, reads_helix_sharded_ctx_pages, shift_label), so a helix-sharded context is shared only with a drafter that reads it. Include cache-manager, peer-metadata and speculative-worker regression coverage. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
…f shared draft KV A DWDP context MPI world starts its context servers through disaggregated_mpi_worker from a config that lists no generation servers; those run as a separate launch. The shared draft KV pre-check required speculative_config in both sections of that config, so every context rank exited with TRTLLM_DISAGG_SHARED_DRAFT_KV=1. Require speculative_config on every listed server instead and leave the cross-side match to the peer handshake, which already rejects mismatched draft layouts. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
xguannv
force-pushed
the
user/xguan/k3-rubin-modeling-v2
branch
4 times, most recently
from
November 10, 2026 13:53
533e317 to
ac90a78
Compare
… supported page sizes A loaded standalone DSpark drafter decides whether it reads the target's KV pages through reads_managed_ctx_pages. The DFlash base answered True for any page size under the TRTLLM backend, so a GQA DSpark drafter with a target at 128-token pages was put in the shared pool and its worker then rejected the page size at initialization. Return True only for the page sizes the trtllm-gen block decode supports, as before the drafter hook: at other sizes the drafter keeps its own pool. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
…ybrid layout With the FP4 MLA layout, the hybrid manager scans and zeroes the KV cache at warmup through _iter_fp4_mla_physical_pool_views, which listed every physical pool except the one holding the appended draft layers. Yield that single-buffer pool too, viewed with the draft layers' dtype and page shape, so check_invalid_values_in_kv_cache(fill_with_zero=True) covers it. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
…ing in the docs; test the reverse dtype mix and the draft bytes in the hybrid estimate The speculative-decoding guide said a Kimi K3 hybrid target keeps a standalone DSpark MLA drafter's layers with the CUTEDSL, TRTLLM or VANILLA drafter backend. Whether a drafter shares is decided by its reads_managed_ctx_pages declaration, and the default declaration covers only the TRTLLM backend at trtllm-gen page sizes, so the guide now says that; the list of separate-manager cases names the helix and opt-in exceptions. The FP8 MLA context test also rejects BF16 K with FP8 V. A CPU test checks that MambaHybridCacheManagerV2.get_cache_size_per_token adds the appended draft layers' bytes per token to the slope and does not pass layer_specs or shared_draft_block_width on to the hybrid estimate. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
An embedded DSpark draft (DeepSeek-V4 / V4.1) is a DSv4DSparkForCausalLM, not a DFlashForCausalLM, and has no reads_managed_ctx_pages. The shared draft KV gate read model.draft_model for every DSpark run and called that hook, so creating the KV cache raised AttributeError where main reported that the drafter has no paged block attention. _loaded_standalone_dspark_drafter returns None for an embedded draft, and _dflash_unify_blocker returns its reason before reading any hook when the drafter has no block attention backend. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
…sided config, and the receiver rejects a draft layout mismatch With TRTLLM_DISAGG_SHARED_DRAFT_KV unset, disaggregated_mpi_worker set it to "auto" for a config that lists only context servers (a DWDP context MPI world), where main set "0". A generation worker launched separately with the variable unset keeps a separate draft cache, so the two sides disagree. _shared_draft_kv_default turns sharing on only when the config lists both sides and passes the check; a one-sided config shares only when the variable is "1". The draft layer count and attention TP checks ran only in the sender's register(), whose errors the listener thread logs, so a receiver whose draft layout disagreed waited out kv_transfer_timeout_ms. They move to PeerRegistrar.validate_draft_layers_compatible, which the receiver also runs before it registers with a sender, so its requests to that sender fail with PeerIncompatibleError. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
…h holds The cap is set after the final KV cache and CUDA graphs, so the 1 GiB kept outside it comes out of the memory left free at that point. With less than 1 GiB free (a high free_gpu_memory_fraction on a small GPU) the cap dropped below memory_reserved(), and with free + reserved under 1 GiB the fraction was negative and set_per_process_memory_fraction raised ValueError. The margin is now min(1 GiB, free memory). Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
Serialize DWDP worker URLs as YAML strings and publish every serving role to the same discovery directory when the launcher removes Slurm variables. This lets the frontend discover context workers and avoids waiting on a different directory. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
…rd at startup With num_serve_frontends > 1, a secondary frontend forwards the disaggregated server's POST /steady_clock_offset to the launcher frontend over a Unix socket. While the launcher frontend is not listening yet the forward is refused: the frontend logs a warning that names ClientConnectorError and ConnectionRefusedError and answers 503, and the disaggregated server leaves that server's perf-metric timestamps unadjusted. The server log scan counted the warning as a server error and failed the run. Treat exactly that warning as benign; a refused forward of any other route, a timeout, or a line that also carries another error keyword is still reported. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
… collectives The kimi_k3_nvfp4__sm_107__default Modeling V2 target's deployment features, the family routing to it, its core layers (linear, MXFP8 GEMM, norm, MLP, embedding, all-reduce, streams), the Lamport collectives and the attention-residual norm, with their ops/ entries and tests. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
The target's MLA attention: context FMHA and split-K context over the FP4 / FP8 KV cache, FP4 MLA decode with the packed reduce, the fused q path and the Helix tail, with their ops/ entries and tests. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
The target's KDA mixer and its prefill, MTP-decode and fused-projection kernels, with their ops/ entries and tests. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
The target's routed and shared experts (router, route tail, small-batch MoE, ULL front, MoE tail) and its low-M / mid-M FP8 and MXFP8 GEMMs with their push variants, with their ops/ entries and tests. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
The target's DSpark drafter and spec-decode tail (vocab-parallel head, Markov chain, acceptance, draft KV), with their ops/ entries and tests. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
The target's model (modeling.py), its checkpoint loading (weights.py) and the runtime modules they use, with the model-level tests. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
…lists GSM8K accuracy of the target's disaggregated deployments, their perf-sanity configurations, and the CI and QA test lists for the target's tests. Both roles of every deployment store the BF16 shared and latent MLP projections as FP8 (KIMI_K3_FP8_WEIGHT_READ_MOE_MLP=1; without it the DWDP context servers run out of memory), keep the KV cache of the last two MLA layers in FP8 (TLLM_KIMI_K3_FP8_KV_LAYERS=91,92) and use expandable PyTorch allocator segments. The context servers cap their token capacity (1048576 with DWDP over 16 ranks, 131072 for the low-latency deployment): GPU 0 of each node also holds the CUDA contexts of the node's other MPI ranks, which leaves the default pool too little room there. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
… and TCP transports The perf-sanity launcher clears UCX_TLS for clusters without a UCX rule, so the context and generation workers of the Kimi K3 SM107 gates ran with every UCX transport enabled. The KV cache pool sized during the startup memory profiling then stays registered with UCX after it is released, and its memory does not come back for the final KV cache allocation: on the dep16 generation server, the KV cache manager's shutdown returns 2.8 of the pool's 35.9 GiB with UCX_TLS unset, and all of it with UCX_TLS limited to cuda_ipc,cuda_copy,sm,self,tcp. Set UCX_TLS=cuda_ipc,cuda_copy,sm,self,tcp in the shared worker environment of the four gate configurations (the transports the GB300 rules in tests/test_common/cluster_env.py pin), and drop the context servers' max_tokens caps, which only made room for the memory that was not returned. Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
xguannv
force-pushed
the
user/xguan/k3-rubin-modeling-v2
branch
from
November 10, 2026 17:30
ac90a78 to
2fc39e6
Compare
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
Kimi K3 NVFP4 on Rubin (SM107) as one Modeling V2 target,
kimi_k3_nvfp4__sm_107__default, plus the runtimechanges it needs, so that it builds and runs on main by itself.
Each layer below is self-contained in its commits; review per layer.
1. Runtime (5 commits)
76637bdd60kvCacheManagerV2Utils.cu): bounded chunks per row, several CTAs per long table; layout unchanged.a1b8638c2d6b20017b37attentionOp.cpp,mlaKernels.cu, TRTLLM backend); BF16 inputs keep the existing path.1921268140TRTLLM_DISAGG_SHARED_DRAFT_KV=1.f0cbfb8be3disaggregated_mpi_workerchecks a one-sided config (a DWDP context MPI world lists no generation servers) for its own side; the peer handshake rejects mismatched draft layouts.2. DWDP test discovery fix (1 commit)
1fc5fa945d: DWDP worker URLs are serialized as YAML strings and every serving role publishes to the same discoverydirectory when the launcher strips Slurm variables, so the disaggregated frontend finds the context workers.
3. The Kimi K3 SM107 target (7 commits)
tensorrt_llm/_torch/_experimental/modeling_v2/models/kimi_k3_vl/kimi_k3_nvfp4__sm_107__default/; the family routerreturns it for
KimiK3ForConditionalGenerationon SM107. One target serves the four deploymentsfeatures.pyknows:
tp1(one-GPU DWDP context server),helix8,helix16_moetp4ep4(ULL) anddep16; the profile is selectedbefore the target's modules are imported.
Under the default
modeling_v2=autothe route selects this target on its own: it claims only SM107, Kimi K3'sexact text shape, the NVFP4 checkpoint, DSpark or no speculative decoding, an NVFP4 or FP8 KV cache, and one of the
four layouts above; anything else raises
NoTargetand the built-in Kimi K3 implementation runs. Like the two KimiK3 SM100 targets it is not
@disable_v2_by_default. The shipped deployment configs setmodeling_v2: require.e175294fc2cd0433ba940b7ec090d194063aaf0b74672aa51d270b61f79cmodeling.py,weights.pyand the runtime modules (moe_runtime,fp8_weight_read,attn_res,text_config)533e3177f4attention_backend: CUTEDSLexplicit, which is whatAUTOresolves to for this MLA drafter), CI / QA test listsThe deployments opt into FP8 storage of the BF16 shared and latent MLP weights (
KIMI_K3_FP8_WEIGHT_READ_MOE_MLP=1,the same opt-in as the built-in Kimi K3 implementation): it is a lossy conversion of checkpoint weights, so the
target's default keeps the checkpoint's BF16; the DWDP context servers need the memory it frees. The ULL deployment
caps the context servers'
max_tokensat 524288 because GPU 0 of each node also holds the CUDA contexts of thatnode's other MPI ranks.
Layout of the target directory:
modeling.py,weights.py,text_config.py,features.py) and one module per block:layers.py,mla.py,kda.py,moe.py,drafter.py,dflash.py, plusmoe_runtime.py,attn_res.py,fp8_weight_read.py.ops/<category>/: theOpWrapperentries (attention, comm, embedding, gemm, moe, norm, spec, ssm); host glue usedby a single entry sits next to it as a private module (
_fp4_mla_decode.py,_mxfp8_gemm.py,_kda_prefill.py,_fla_index.py,_dspark_mla_rope.py).kernels/: the CuTe DSL / Triton kernels (kda/,mla/,moe/,spec/and shared kernel files).modeling.pyimports the block modules (MLA, drafter) eagerly, afterfeatureshas selected the deployment, soan invalid
TRTLLM_K3_ULL3_MLA_GATE_ATfails at import.How the target follows the Modeling V2 rules:
OpWrapperentry inops/<category>/<entry>.py(attention, comm, embedding, gemm, moe, norm, spec, ssm) that calls the CuTe DSL /Triton launcher directly, with a plain-torch
reference,is_valid, andCELLS(the shapes K3 runs, each withits deployment). Tests drive every cell under
validating.(e.g.
kernels/mla/mla_kv_fp8.py, the drafter's split-KV planner, the Helix post-process); the rest go through catalogentries.
modeling.py/weights.pyinstead of one 6k-line file; classes that shared a name with aModeling V1 class are renamed (
KimiK3DefaultForCausalLM,K3MoERuntime,K3MLADSparkForCausalLM, ...).Known gaps (listed so review can decide):
k3_fp4_mla_decodeandfp4_mla_kv_append(a reference needs the KV cache manager'sring / packed-V layouts), the three drafter all-reduce sites (
k3_draft_ar_norm,k3_draft_embed,k3_draft_gemm_push; their 2-rank test compares them bitwise with the C++ fused all-reduce), andk3_ull_front,k3_moe_tail1,k3_shared_persist.bmm_out/cublas_mmdeclare ARCHS {SM_100, SM_103}; the target calls them on SM107.kimi_k3_vlowner line.Test Coverage
Unit tests (
tests/unittest/_torch/modeling_v2/models/kimi_k3_vl/kimi_k3_nvfp4__sm_107__default/, listed inl0_gr100; multi-rank ones inl0_gb300_multi_gpus; CPU ones inl0_cpu):test_modeling_v2_contracts.py+test_modeling_v2_routing.py+test_disagg_utils.py1512 passed / 260 skipped / 0 failed (1701850; skips areSM107-only kernels); the K3
l0_cpuentries 45 passed / 48 skipped / 0 failed (1702502).346 passed / 70 skipped (other models' sm_103-only targets). The 8 skips are the context FMHA instrumentation
checks, which need the internal CuTe DSL (
cute.arch.nanosleep). The 2 failures were in the plain-torchreferenceoffp4_mla_gather_mxfp8(a local reused the name of the context-offset input when the token countis not a multiple of 128), fixed in this head; on this head both cells and the rest of
ops/attentionpass(28 / 0, hecate 1020360).
Runtime layers 1–2 (GB300, built on the previous main base): the 14 test files they add or change pass (the cache
transceiver file 269 / 0 with UCX bound to the node's Ethernet NIC); the existing suites of the touched modules
(KV cache, speculative, disaggregated) show no new failures (3
test_kv_cache_v2_vocab_sizeqwen3_5_hybrid cases failidentically on main); C++
mlaPreprocessTest6 / 6 andkvCacheManagerV2CudaContextScopeTest3 / 3. Their 54SM107-only cases (
test_fp4_mla.py,test_fp4_mla_kv_transfer.py,test_kimi_k3_mla.py,test_dspark_cute_dsl_attention.py,test_kimi_kda_fused_verify_parity.py) on Rubin (hecate 1020178): 219 passed,2 failed. The 2 are
test_fp4_mla.py::test_fp4_mla_helix_scatter_uses_global_rope_and_local_cache_position[active/inactive]; that test and the code it covers are unchanged from main (it compares the 128-row padded Q
buffer with the 4 active rows), so it fails on SM107 without this PR as well.
Bitwise identity with the pre-refactor model (same kernels, same launch order; GB300):
helix16_moetp4ep4 profiles and at TP4 over MNNVL: OpWrapper conversion and module split each bit-identical to the
pre-refactor code (jobs 1813764), MoE layer (M = 1..1024, tp1 / adp1 / TP4) and DSpark drafter (tp1, dep16, TP4,
helix8, helix16) bit-identical (1814540).
model (tp1, dep16, helix8, helix16_moetp4ep4, TP4) and MoE layer + drafter bit-identical (ocijhb 1698165).
Same build, final code against the pre-refactor code: whole model, MoE layer and drafter bit-identical in every
profile and on every rank (ociaga 1815190).
profile and at TP4, MoE layer and drafter bit-identical (ocijhb 1701855).
Accuracy (Rubin, GSM8K strict-match, disaggregated, the deployments in
test_accuracy.py): 95.91 (helix8 bs20DSpark5), 96.13 (dep16), 96.29 (ULL helix16 MoE TP4 x EP4) on the pre-refactor code; this head: TODO (hecate 1021419 / 1021420 / 1021421).
Performance (Rubin, same-allocation ABBA against the previous Rubin branch, three points c383 / c240 / c24):
TODO. The compared build has the same runtime as this PR; one known difference between the arms besides this code:
the previous branch also has a NIXL request-release gate (#21612, not part of this PR) that the compared build and
this PR do not; all three points use the NIXL Python transceiver.
PR Checklist